合并受阻
变更摘要
此 PR 修复了 NPU inductor 在元素总数超过 2^31 时因 int32 索引溢出导致的问题。先前将所有 arange 切片统一提升为 int64 会导致 UB 翻倍并引发挂起（bug 507034）；若全部保持 int32 则线性地址会回绕造成错误（bug 507035）。该变更采用"变体 C"方案：保持大规模向量切片为 int32，仅在指针运算的线性地址合成处，将那些绝对值 × (轴长度-1) 超过 2^31 的溢出加数项（如 268435456*x1）提升为 int64，同时将掩码比较的 LHS 在静态可判定安全时回退为 int32 以保持在向量单元上执行。
主要改动
-
恢复
dtype_to_str中的 int64 索引类型：在NPUTritonKernel.dtype_to_str中新增对torch.int64的处理，返回"tl.int64"，覆盖上游补丁将 int64 退化为 int32 的行为，确保超过 2^31 元素的内核正确触发 int64 索引路径。 -
新增
_npu_promoted_overflow_terms方法：从 sympy 表达式中识别哪些乘加项（c * axis）在 int32 下会溢出，贪心地按溢出量降序选出需要提升的项，确保剩余未提升的部分一定在 int32 范围内。 -
在
index_to_str中实现按使用点提升：当index_dtype为"tl.int64"时，对表达式文本中的溢出加数项追加.to(tl.int64)转换（如268435456*x1→268435456*x1.to(tl.int64)），使提升仅作用于小规模/标量潜伏轴切片，而非大规模 arange 切片。 -
arange 切片保持 int32：移除
codegen_range_tree、_indexing_range_code和iteration_ranges_scalar_code中对 arange 切片和标量偏移的 int64 上溯转换，统一硬编码为tl.int32，由使用点的选择性提升覆盖溢出风险。 -
_mask_cmp_lhs增强 int32 回退逻辑：新增numel和index_dtype参数，当索引类型为 int64 但轴长度静态已知且小于 2^31 时，将比较左操作数显式转换回tl.int32，使比较保持在向量单元执行，避免因缺少原生 int64 向量比较而降级为标量循环。
✅ 跳过 docs ci 检查，没有需要检查的文档文件
fix(npu-inductor): int32 index overflow — 构造性正确升宽（int64 地址 + int32 tile）
【合入来源】
【问题背景】
总元素数超过 2^31 的 kernel，其线性地址 = Σ(stride × 轴索引) 会越过 int32 表示范围，形成两条死路：
268435456*x1（x1≥8）回卷成负数 →ptr + 负偏移→ 507035（MTE 非法访存，vector core exception）硬件事实：AIV 的 UB 预算有限；int64 向量 ALU 为仿真（compute-bound 实测 3.5×）。二者把「全 int32」与「全 int64」同时判死，正确解必须让 i64 精确出现在最小必要位置集合上。
【方案设计】
核心思想：构造性正确（Constructive Correctity）
正确性由 triton 类型提升规则构造性保证：每个地址算术至少含一个 i64 操作数 ⇒ 任意形状、任意项形态、任意求值顺序下不可能回卷。不依赖任何求界分析、形状快照或运行时 guard。
实拍生成形态（8388609×256 sum，恰 2^31+256 元素）：
tmp0 = tl.load(in_ptr0 + (256*x0.to(tl.int64) + r0_1.to(tl.int64)), r0_mask & x0mask, ...) tl.store(out_ptr0 + (x0.to(tl.int64)), tmp2, x0mask) # tl.arange / tl.full 全零 int64 —— tile 保持 int32（逐行断言钉死）机制分层
select_index_dtype()（上游继承）+_npu_should_widen_addressModularIndexing/ 轴→len−1 代换后整表达式 bound 超 int32）；所有平局向 fail-safe 破，换形状经普通 guard 重编译重判dtype_to_str解除 int64 降级force_linearizepid.to(i64)*XBLOCK混合广播升宽全部 header tile）；检测异常时 warn + 安全侧 TrueNpuWiden+ printer hook.to(tl.int64)（ToFloat 同款惯用法）；节点级 subs、零文本匹配——负系数项、复合 stride、FloorDiv/Mod 项全部正确（正则方案对负项渲染形态(-c)*x从不命中）equal_to特化 uint32→i64 vcast 被 BiShengIR 拒绝）；一律别名 i64 形参，整条 odometer/rsplit 标量链经提升规则自然 i64；动态兄弟多轴角落 fail-loudnpu_triton_compute_type/npu_triton_store_type/value_expr补丁_mask_cmp_lhs收窄 + 终端性防护mask_cmp_fp32（默认关）开启时的 fp32 比较路径在 triton-ascend 存在已知 lowering 缺陷（507034 挂死），终端性保证收窄覆盖的 kernel 无论开关状态都结构上不可达该路径。行为由纯文本断言钉死：默认关闭时四种形态布局逐字节不变；开启时 int64-narrow 形态保持 int32 比较ceil(xnumel/XBLOCK)，防过读 507035 / 静默丢 tile_load_cached_autotuning漏 pop 上游恒写的found_by_coordesc/triton_cache_hash→ 任何后端热缓存重跑全 config 失败；补 pop + TE 边界 constexpr 过滤mask 比较的防护机理：mask 不参与地址算术，走独立的比较路径。
_mask_cmp_lhs对 int64 索引 kernel（静态轴长 <2^31）生成(index).to(tl.int32) < numel——int32 比较落在 AIV 的确定性路径上；该分支return早于mask_cmp_fp32的 fp32 cast 判断，因此开关开启时控制流也止步于 int32 比较，不会生成to(tl.float32)形态。测试以两态断言钉死：默认关闭时布局逐字节不变（零行为差异）；开启时 int64-narrow 形态必须保持 int32 比较。设计权衡（为什么不是其它方案）
arange/full/pid全.to(i64)）：UB 翻倍是功能故障（507034）非性能问题；to_dtype节点，地址算术是 sympy 渲染，层位不可见；【改动文件】
codegen/triton.pycodegen/npu_header.py_npu_emit_axis_numel（≥2^31 别名 i64 形参）、blocks 守卫与说明npu_triton_heuristics.pyruntime/autotune_cache.pytest/_inductor/test_triton_experimental_int32_overflow.py【自测信息】
tl.int64（逐行断言）防回归钉子五类：with_index 哨兵/类型配对、in-range 零 int64、类型路由策略、上游类型助手调用面 28 位点快照（增改必红，torch 升级触发重审计）、≥2^31 裸字面量 NotRegex。大张量用例均带
@skipIfInsufficientHBM保护。