Menu
Triton Ascend 调试排查清单和常见错误速查表，包括编译错误、运行时错误、精度问题和性能问题的诊断方法。适用于内核代码生成、出现错误需要定位原因、或需要验证代码正确性的调试场景
Installer avec Codex ou Claude Copiez ce prompt, collez-le dans Codex, Claude ou un autre assistant, puis laissez-le vérifier la page du skill et l'installer pour vous.
Basé sur la classification professionnelle SOC
矩阵乘法矩阵乘法 A[M, K] @ B[K, N] = C[M, N]中，大K维度矩阵乘法(K>>M,N)优化：针对M/N较小但K极大(如M=N=256,K=131072)的场景，Split-K切分K维度并行化、Workspace+Reduce替代全局同步，实现显著性能提升
Triton Ascend hard API restrictions and forbidden syntax. MUST-follow rules that apply to every kernel: forbidden control flow (return/break/continue/lambda/while), tensor slice/index restrictions, scalar conversion rules, BLOCK_SIZE upper bound. Violating any of these produces a compile or runtime error on Ascend.
Triton Ascend 性能优化通用策略: BLOCK_SIZE 选择 (1024-2048 for elementwise, must be <65536), grid configuration (use VEC_CORE_NUM / CUBE_CORE_NUM, 2D/3D grid for matmul / conv / reduce, 1D grid + inner loop for elementwise / pointwise), 256B alignment for memory transfers, autotune block-size patterns, fp16 / fp32 precision conversion. Bind via keywords like matmul, elementwise, reduce, block_size, grid, autotune, alignment, fp16, fp32, tile, interleaved-loop, cube-core, vec-core.
通过 adaptive_search 或 evolve 搜索式 workflow 生成优化算子。 后台 silent mode 执行，轮询监控进度。
适用于归约(reduce)类算子和含归约子步骤的复合算子（如归一化）的优化指南。典型算子包括：sum, mean, max, min, prod, argmax, argmin, cumsum, cumprod, softmax, logsoftmax, layernorm, rmsnorm, groupnorm, instancenorm, batchnorm, l1norm, l2norm, frobeniusnorm, var, std, average_pooling, sum_pooling 等。特别重要：当归约维度不是最后一维（如 dim=1 归约 shape=[B,F,D1,D2]），需要正确处理多维索引和两阶段归约。包含 PyTorch normalized_shape 多轴归一化语义说明。不适用于纯逐元素运算或矩阵乘法。如果算子是损失函数（先逐元素计算再全局归约），应选择 elementwise-reduce-fused 指南。
CPU C++ 算子核心概念、标准结构模式、KernelBench 代码规范和内嵌扩展方法
| name | triton-ascend-debugging |
| description | Triton Ascend 调试排查清单和常见错误速查表，包括编译错误、运行时错误、精度问题和性能问题的诊断方法。适用于内核代码生成、出现错误需要定位原因、或需要验证代码正确性的调试场景 |
| category | fundamental |
| version | 1.0.0 |
| metadata | {"backend":"ascend","dsl":"triton_ascend","hardware":"Atlas A2, Atlas A3"} |
tl.constexpr 是否只在内核参数中使用？for i in range(pid, total, core_num)？tl.atomic_add 等）？b[0]、b[i:j]）？tl.arange生成的张量误用了tl.get_element？tl.get_element、tl.extract_slice等）？| 禁止写法 | 替代方案 |
|---|---|
return / break / continue | 使用 mask 控制流程 |
| lambda 表达式 | 内联函数或 tl.where |
链式布尔运算 a and b | 分步计算 mask：m1 = ...; m2 = ...; m = m1 & m2 |
张量直接索引 tensor[i] | tl.load(ptr + offset) / tl.store(ptr + offset, val) |
Python 切片 b[0] / b[i:j] | tl.get_element / tl.extract_slice / tl.insert_slice |
对 tl.arange 结果用 get_element | 直接计算索引值 |
while 循环 | for i in range(MAX): if i < n: |
range() 混用运行时变量和 constexpr | 全 constexpr 的 range(0, N, BLOCK_K) + 循环体内运行时 if |
tl.float16(scalar) | scalar.to(tl.float16) |
tl.constexpr 在 host 侧使用 | 仅在 kernel 参数中使用 |
| if-else 中负偏移 | tl.maximum(offset, 0) |
复杂 tl.where 用于内存偏移 | 拆分为 if-else 静态分支 |
| 错误类型 | 典型症状 | 常见原因 | 解决方案 |
|---|---|---|---|
| UB/CBUF 溢出 | ub overflow, requires X bits while 1572864 bits available | BLOCK 尺寸过大或中间变量过多 | 缩小 BLOCK 尺寸；减少同时活跃的 tensor 数 |
| HiVM vsel 错误 | hivm.hir.vsel: Unsupported op for finding the root alloc | 嵌套 mask + tl.where 组合过于复杂 | 用乘法替代 tl.where：a * mask.to(dtype) |
| 内存越界访问 | 运行时错误、结果异常、随机崩溃 | load/store缺少mask或boundary_check | 添加正确的mask或boundary_check保护 |
| Grid超限 | 编译失败或运行时错误 | grid总大小超过65535 | 使用交错循环for i in range(pid, total, core_num)或连续分块处理 |
| 控制流错误 | unsupported AST node type: Continue | 使用了return/break/continue | 改用 if-else 包裹逻辑 |
| while循环错误 | 编译失败（Ascend后端） | 使用了while循环 | 改用for + if替代：for i in range(MAX): if i < n: |
| constexpr 索引 | ValueError('unsupported tensor index: constexpr[0]') | 对 tl.sum 等返回的标量做 [0] 索引 | 直接使用标量结果，不要索引 |
| 切片语法错误 | 编译失败 | 使用了b[0]或b[i:j]直接切片 | 使用tl.get_element或tl.extract_slice |
| tl.arange索引错误 | 编译失败 | 对tl.arange结果使用get_element | 直接计算索引值而非提取 |
| 类型转换错误 | cast incompatible | 隐式 cast 或使用tl.float16(scalar) | 用 .to(tl.float16) 或显式指定 acc dtype |
| constexpr误用 | 编译失败 | 在host侧使用tl.constexpr | 仅在kernel参数中使用tl.constexpr |
| Stride设置错误 | 计算结果错误、数据错位 | stride参数计算或传递错误 | 验证stride设置，检查tensor.stride() |
| 数值不稳定 | 结果为NaN或Inf | softmax/sqrt等操作溢出 | 减去最大值、检查非负、使用float32 |
| 数据竞争 | 结果不确定、每次运行不同 | 多program并发写入同一位置 | 使用tl.atomic_add等原子操作 |
| BLOCK_SIZE过大 | 编译失败或运行时错误 | BLOCK_SIZE超过65536或硬件限制 | 减小BLOCK_SIZE，使用循环处理 |
| tl.where偏移计算 | 编译失败（Ascend后端） | 在内存偏移中使用tl.where | 改用if-else静态分支处理 |
| 性能低下 | 运行缓慢 | 内存访问不连续、切分不合理 | 优化内存布局、调整BLOCK_SIZE、使用block_ptr |
| 运行时range边界崩溃 | bishengIR crash | range()的start/stop混用运行时变量和constexpr | 改用全constexpr的range(0, N, BLOCK_K)，循环体内用运行时if跳过 |