القائمة
Triton Ascend 失败后的调试定位清单和常见错误速查表。仅在验证、编译或运行失败后使用，用于帮助 Conductor 解释报错、定位根因并给出修复建议。生成期硬约束请参考 api-rules、hardware-constraints、memory、grid-config 和对应 guide。
التثبيت باستخدام Codex أو Claude انسخ هذا Prompt والصقه في Codex أو Claude أو مساعد آخر ليراجع صفحة Skill ويثبّتها لك.
استنادا إلى تصنيف SOC المهني
AscendC direct-invoke 崩溃/挂起修复索引：Kernel timeout、hang、Segmentation Fault、aic error、buffer 死锁、plog/memcheck 调试。
AscendC direct-invoke 精度失败修复索引：输出全 0/随机值、DataCopy 对齐、EnQue/DeQue 同步、FP16/FP32 精度、Cast RoundMode、DumpTensor 分段定位。
AscendC direct-invoke 工程契约：WA 使用 kernel.py + ascendc_op/，ModelNew 调用 torch.ops.npu.*，adapter 负责复制工程、CMake 构建和 npu-arch patch。适用于 dsl=ascendc 的 autoresearch 任务。
把注册式 AscendC 算子迁移为 direct-invoke 工程的保真原则：kernel 算法和 tiling 公式不乱改，只替换注册框架胶水、入口 ABI、host launch 与 PyTorch extension。
CATLASS TileShape 与 on-chip 缓存容量约束：L1/L0A/L0B/L0C 预算公式、fp16/fp32 Pingpong 双缓冲、512B 对齐与排布对 Tile 选型的影响。调参前必读。
CATLASS Gemm 性能调优：DispatchPolicy、Tile 与分核负载均衡、Swizzle、何时用 padding/Split-K/Preload。面向 AR 修改 catlass_kernel.asc 中的类型别名。
| name | triton-ascend-debugging |
| description | Triton Ascend 失败后的调试定位清单和常见错误速查表。仅在验证、编译或运行失败后使用，用于帮助 Conductor 解释报错、定位根因并给出修复建议。生成期硬约束请参考 api-rules、hardware-constraints、memory、grid-config 和对应 guide。 |
| category | fix |
| version | 1.0.0 |
| metadata | {"case_type":"fix","backend":"ascend","dsl":"triton_ascend","hardware":"Atlas A2, Atlas A3, Atlas A5"} |
本 skill 只用于失败后的诊断和修复建议，不作为初始生成约束注入 KernelGen。 如果结论属于“生成时必须遵守”的规则，应落回 fundamental 或对应 guide。
tl.load / tl.store 是否都有 mask 或 boundary_check？tensor.stride() 一致？ptr + offsets，输入是否先 .contiguous()？tl.where / mask 组合过深？tl.constexpr 是否只在内核参数中使用？for i in range(pid, total, core_num)？tl.atomic_add 等）？b[0]、b[i:j]）？tl.arange生成的张量误用了tl.get_element？tl.sum / tl.max 这类 reduction scalar 结果做了 [0] 索引？tl.get_element、tl.extract_slice等）？| 错误类型 | 典型症状 | 常见原因 | 解决方案 |
|---|---|---|---|
| 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跳过 |
| 错误位置分布 | 优先怀疑 | 修复方向 |
|---|---|---|
只集中在最后一维尾部，例如 dim2: [16:17] | tail mask / boundary_check 漏写 | 检查最后一维 load/store mask，block_ptr 加 boundary_check=(0, 1) |
| 某一整行或整列错误 | 行列 offset、stride、转置或 sub_vec_id 偏移错误 | 复核多维 index 分解和 block_ptr shape/strides/offsets/order |
| 所有 tile 都像最后一个 tile | A5 亲和单 buffer 被覆盖 | cube 每个 tile fixpipe 后 wait vector 释放，vector 读完后 set buffer-free |
| 误差随 K 增大明显变大 | fp16 长链累加精度损失 | fp32 accumulator；必要时参考 Kahan precision fix |