triton-ascend-error-fix

1. UB / CBUF 溢出

• 报错特征: cbuf overflow
• 根因: 分块参数（BLOCK_M/N/K）过大，超出 Ascend UB 容量
• 修复: 减小分块参数，通常 BLOCK_M=64, BLOCK_N=128 是安全起点

# 错误：分块过大导致 UB overflow
BLOCK_M, BLOCK_N, BLOCK_K = 128, 256, 256

# 修复：安全起始点
# CUBE (matmul fp16): BLOCK_M=64, BLOCK_N=64, BLOCK_K=32
# CUBE (matmul fp32): BLOCK_M=32, BLOCK_N=32, BLOCK_K=32
# VEC (elementwise):  BLOCK_SIZE=1024~2048
BLOCK_M, BLOCK_N, BLOCK_K = 64, 64, 32

4D tensor 矩阵乘法中 batch 维度额外占用 UB 空间，建议将 batch 展开到 grid 维度而非内层循环。

2. BiShengIR / HiVM 编译失败

• 报错特征: hivm.hir.vsel: Unsupported op for finding the root alloc、Failed to run BiShengHIR pipeline
• 根因: 编译器后端不支持复杂的 mask 组合或指针运算模式

2a. 内联地址计算过于复杂

# 错误
tl.store(c_ptr + off_m[:, None] * stride_cm + off_n[None, :] * stride_cn, acc, mask=mask)

# 修复：拆分为中间变量
c_ptrs = c_ptr + off_m[:, None] * stride_cm + off_n[None, :] * stride_cn
tl.store(c_ptrs, acc, mask=mask)

2b. tl.where + 复杂 mask 导致 vsel 错误

# 错误：嵌套 mask + tl.where 触发 hivm.hir.vsel 错误
a_tri_mask = a_offsets_k[None, :] >= a_offsets_m[:, None]
a_valid_mask = a_mask_m & a_mask_k
a = tl.where(a_tri_mask & a_valid_mask, a, 0.0)

# 修复：改用乘法替代 tl.where（将 bool mask 转为 float 后相乘）
a_tri_mask = (a_offsets_k[None, :] >= a_offsets_m[:, None]).to(tl.float16)
a_valid_mask = (a_mask_m).to(tl.float16) * (a_mask_k).to(tl.float16)
a = a * a_tri_mask * a_valid_mask

3. Triton 语法限制违反

3a. 禁止 continue / break / return

# 错误：unsupported AST node type: Continue
for i in range(N):
    if condition:
        continue
    do_work()

# 修复：用 if-else 包裹
for i in range(N):
    if not condition:
        do_work()

3b. constexpr 索引错误

# 错误：ValueError('unsupported tensor index: constexpr[0]')
result = tl.sum(data, axis=0)
tl.atomic_add(out_ptr, result[0])

# 修复：tl.sum 已返回标量，直接使用
result = tl.sum(data, axis=0)
tl.atomic_add(out_ptr, result)

3c. tensor.cast 类型不兼容

# 错误：cast incompatible shapes
result = tl.dot(a_fp16, b_fp16)

# 修复：显式指定 fp32 累加器
result = tl.dot(a_fp16, b_fp16, acc=tl.zeros([M, N], dtype=tl.float32))