triton-ascend-ascend-hardware-constraints

Ascend 硬件约束与编译器限制

不同型号的硬件具体容量不同，算子生成时会同步传入硬件信息文档，以下公式中的容量值请参考该文档。

1. 存储层级预算

CUBE 路径（matmul / tl.dot）

Matmul 数据走 L0A/L0B/L0C，不经过 UB：

缓冲区	用途	约束公式
L0A	左矩阵 A tile (m0 × k0)	m0 × k0 × sizeof(A.dtype) ≤ L0A容量
L0B	右矩阵 B tile (k0 × n0)	k0 × n0 × sizeof(B.dtype) ≤ L0B容量
L0C	结果 C tile (m0 × n0)，支持累加	m0 × n0 × sizeof(C.dtype) ≤ L0C容量

计算示例（以某硬件 L0A = 64KB 为例）：

• fp16（2 字节/元素）：可容纳 32K 个元素 → BLOCK_M=128, BLOCK_K=256 恰好填满
• fp32（4 字节/元素）：可容纳 16K 个元素 → BLOCK_M=128, BLOCK_K=128 恰好填满

选择 tile 尺寸时，确保三个缓冲区都不超限。fp32 占用是 fp16 的 2 倍，需相应缩小 tile。

VEC 路径（element-wise / reduce / norm）

向量运算数据走 UB：

缓冲区	用途	约束公式
UB	所有活跃 tensor 和中间变量	BLOCK_SIZE × sizeof(dtype) × 活跃tensor数 × multi_buffer系数 ≤ UB容量

编译器启用 auto-multi-buffer 后，实际占用约为基础量的 2~3 倍。kernel 中的中间变量（如 tl.where 产生的临时缓冲）也占用 UB，实际占用会显著高于 BLOCK_SIZE × sizeof(dtype) × 输入数。

tile 选择策略：从较大 BLOCK_SIZE 开始尝试，遇到 ub overflow 编译错误时逐级缩小。

2. bishengIR 编译器已知限制

2.1 range() 边界不可混用运行时变量

# 编译器崩溃（bishengIR SIGABRT）
for k in range(start_n, start_m + BLOCK, BLOCK_K):
    ...

start_n、start_m 是运行时值，BLOCK、BLOCK_K 是 tl.constexpr。这种混合用法会导致编译器内部错误。

规避方案：使用全 constexpr 的 range，在循环体内用运行时 if 跳过无效迭代：

for k in range(0, N, BLOCK_K):  # N 和 BLOCK_K 都是 constexpr
    # 可选：运行时条件跳过无效块
    ...

2.2 复杂 mask + tl.where 导致 HiVM 错误

当嵌套 mask 组合传入 tl.where 时，编译器后端可能报 hivm.hir.vsel: Unsupported op for finding the root alloc。

规避方案：用乘法替代 tl.where，将 bool mask 转为 float 后与数据相乘：

# 触发 hivm.hir.vsel 错误
a = tl.where(tri_mask & bounds_mask, a, 0.0)

# 规避：mask 转 float 后相乘
a = a * tri_mask.to(tl.float16) * bounds_mask.to(tl.float16)

2.3 其他编译器限制

详见 debugging 文档中的「禁止使用的语法」完整列表。

3. Strided memory access 的性能代价

Ascend 硬件对非连续内存访问有显著性能惩罚。当 kernel 的核心路径包含 stride > 1 的内存访问模式（如 pooling 的滑窗、dilated convolution 的间隔采样），Triton 生成的代码需要逐元素或小块 gather，而 CANN 原生算子可能使用硬件数据搬运单元（MTE）的专用模式，性能差距可达数十倍。